/* 
PCLP Solvers.h
----------------------------------------------------------------------
DESCRIPTION:
	Header contains wrapper functions for several implementations of LP solver for Packing Covering problems.
   	  

*/


#ifndef GK_PC_SOLVER_CU_H
#define GK_PC_SOLVER_CU_H 




#include <DataTypes.h>

#ifndef LP_STATUS
#define LP_STATUS
typedef enum _LP_status {LP_FEAS,LP_INF} LP_status;
#endif







/* ************************************************************************************************  *
                                       CUDA GKPC Solver 
*  ************************************************************************************************* */


template<class T>
long countNZ(T *ptr, size_t size)
{
	long nz = 0;
	for(index i=0;i<size;i++)
			nz += ptr[i]!=0 ? 1 : 0;
	
	
	return nz;
}


// ********************************** DEVICE error checking ***********************************
void CHECK_CUDA_ERROR(cudaError_t error, char *name);



/* ********************************** device functions ******************************** */

__device__ dreal atomicExch(dreal* address, dreal val);


template <class myType>
__device__ void warpReduce(volatile myType* data);


template <class myType>
__device__ void warpReduceMin(volatile myType* data, volatile  index* data_idx, index tid);

template <class myType>
__device__ myType reduceSum(volatile myType* data);

template <class myType>
__device__ void reduceMin(volatile myType* data,myType *min,volatile index *data_idx,index *min_idx);


/* ********************************** kernels ****************************************** */

// function prototypes 
__global__   void cuCalculateP(matrix_type *A,real *b,dreal *P,dreal *x, dreal *y, dreal delta, integer M, integer N);
__global__   void cuCalculateR(matrix_type *A, real *c, dreal *y,dreal *R,index *R_idx, dreal *rho, index *q, integer M,integer N);
__global__   void cuFindRho(dreal *R, index *R_idx,index *q,dreal *rho, integer N);
 __global__  void cuUpdatePrimalDual(matrix_type *A, real *b, real *c, dreal *x,dreal *y, dreal *P,dreal *R,  
                         index *min_idx, dreal *primal,dreal *dual, dreal *zcDual,dreal *rho,dreal eps, dreal delta, integer M);
__global__   void cuScalePrimalDual(dreal *x, dreal *y, dreal primalScale, dreal dualScale, integer N);

LP_status cuGKPCSol(matrix_type *A, real *b, real *c, dreal *x, dreal *y, integer m, integer n, dreal omega);






#endif

